// autodim.h

// Copyright (C) 2010  traxex

// Author: 	traxex <traxex@yahoo.cn>
// Created: 	11 Aug 2010
// Version: 	1.0
#pragma GCC system_header
#pragma once

namespace traxex{
namespace cuda{

struct __align_as__(int) ArrayDim{
    uint	tx:8;
    uint	bx:12;
    uint	by:12;
    ArrayDim( uint len){
        uint	size = len-1;	//len->(1...len)  size->(0...size)
        *this = *(ArrayDim*)&size;
        if( by > 0)	    bx = (1<<12)-1;
        if( bx > 0)	    tx = (1<<8)-1;
    }
};


struct __align_as__(int) MatrixDim{
    struct{
        uint	t:4;
        uint	b:12;
    }x,y;
    
    MatrixDim( uint ylen, uint xlen){
        uint	ysize = ylen-1;
        y = *(typeof(x)*)&ysize;
        if( y.b > 0)	 y.t = (1<<4)-1;

        uint	xsize = xlen-1;
        x = *(typeof(x)*)&xsize;
        if( x.b > 0)	 x.t = (1<<4)-1;
    }
};

struct CudaDim{
    dim3	block;
    dim3	thread;
  public:
    CudaDim( uint len){				//array
        ArrayDim ad(len);
        block	= dim3( ad.bx+1, ad.by+1, 1);
        thread	= dim3( ad.tx+1, 1, 1);
    }
    CudaDim( uint ylen, uint xlen){		//matrix
        MatrixDim md(ylen,xlen);
        block	= dim3( md.x.b+1, md.y.b+1);
        thread	= dim3( md.x.t+1, md.y.t+1);
    }
};
}
}

#define cusize(n...)	CudaDim(n).block,CudaDim(n).thread

#define	_b		(blockIdx.x)
#define	_bx		(blockIdx.x)
#define	_by		(blockIdx.y)
#define	_bz		(blockIdx.z)

#define _t		(threadIdx.x)
#define _tx		(threadIdx.x)
#define _ty		(threadIdx.y)
#define _tz		(threadIdx.z)

#define _vi		((_by<<20)|(_bx<<8)|(_tx))
#define _mx		((_bx<<4)|(_tx))
#define _my		((_by<<4)|(_ty))
#define _myx		((_by<<4)|(_tx))
#define _mxy		((_bx<<4)|(_ty))


#define _ntx		(blockDim.x)
#define _nty		(blockDim.y)
#define _ntz		(blockDim.z)

#define _nbx		(gridDim.x)
#define _nby		(gridDim.y)
#define _nbz		(gridDim.z)

#define _nt		(_ntrdx*_ntrdy*_ntrdz)
#define _nb		(_nblkx*_nblky*_nblkz)
